 ///////////////////////////////////////////////////////////////////////////////
//! @file	: prefix_sum_line.cl
//! @date   : 2016/03/04
//! @author: guyadong
//! @brief  : Calculates the integral sum scan of an image
////////////////////////////////////////////////////////////////////////////////
#include "common_types.h"
#ifndef CL_DEVICE_LOCAL_MEM_SIZE
#error not defined CL_DEVICE_LOCAL_MEM_SIZE by complier with options -D
#endif
#ifndef SRC_TYPE 
#error not defined SRC_TYPE by complier with options -D
#endif
#ifndef DST_TYPE 
#error not defined DST_TYPE by complier with options -D
#endif
#ifndef INTEG_TYPE 
#error not defined INTEG_TYPE by complier with options -D
#endif
 
#define LOCAL_BUFFER_SIZE (CL_DEVICE_LOCAL_MEM_SIZE/sizeof(DST_TYPE))

#define _KERNEL_NAME(s,d,t) prefix_sum_line_##s##_##d##_##t
#define KERNEL_NAME(s,d,t) _KERNEL_NAME(s,d,t)

#if INTEG_TYPE == INTEG_SQUARE
#define compute_src(src) ((DST_TYPE)src*(DST_TYPE)src)
#define _kernel_name_ KERNEL_NAME(SRC_TYPE,DST_TYPE,integ_square)
#elif INTEG_TYPE == INTEG_COUNT
#define compute_src(src) (DST_TYPE)(src?1:0)
#define _kernel_name_ KERNEL_NAME(SRC_TYPE,DST_TYPE,integ_count)
#elif INTEG_TYPE == INTEG_DEFAULT
#define compute_src(src) (DST_TYPE)src
#define _kernel_name_ KERNEL_NAME(SRC_TYPE,DST_TYPE,integ_default)
#else
#error unknow INTEG_TYPE by complier with options -D
#endif

///////////////////////////////////////////////////////////////////////////////
//! @brief  :	Calculates the prefix sum for each line of an image
// 					Calculates the prefix sum of is_square , 
////////////////////////////////////////////////////////////////////////////////
 __kernel void _kernel_name_( __global SRC_TYPE *sourceImage, __global DST_TYPE * dest, int width, int src_width_step, int dst_width_step){
 	__local DST_TYPE local_block[ LOCAL_BUFFER_SIZE ];
 	__global SRC_TYPE * const src_ptr	= get_global_id(1)*src_width_step + sourceImage;
 	__global DST_TYPE * const dst_ptr	= get_global_id(1)*dst_width_step + dest; 	
 	__global SRC_TYPE * block_src_ptr	= src_ptr;	
 	__global DST_TYPE * block_dst_ptr	= dst_ptr; 	
 	int block_size = 0;
 	DST_TYPE last_sum=0;
	for( int start_x = 0 ; start_x < width ; 
			start_x 				+= LOCAL_BUFFER_SIZE,
			block_src_ptr		+= LOCAL_BUFFER_SIZE,
			block_dst_ptr 	+= LOCAL_BUFFER_SIZE,
			last_sum 			+= local_block[block_size -1] ){
		block_size = min( (int)LOCAL_BUFFER_SIZE, width - start_x );
		// compute prefix sum of a block with local memory		
		local_block[0] = last_sum +compute_src(block_src_ptr[0]) ;
		for( int i=1; i<block_size; ++i)	local_block[i]=compute_src(block_src_ptr[i])+local_block[i-1];
		// copy local_block to dest
		for(int i = 0 ; i < block_size; ++i){
			block_dst_ptr[i]=local_block[i];
		}		
	}
 }
 